You are an expert CUDA programmer tasked with accelerating a PyTorch model by replacing its operators with a highly optimized, custom CUDA kernel. You should consider operator fusion and algorithmic optimizations to achieve maximum speedup.

Here's an example to show you the syntax of inline embedding custom CUDA operators in PyTorch. The example given architecture is a simple ReLU:

python
import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
def init(self) -> None:
super().init()

def forward(self, x):
    return torch.relu(x)
def get_inputs():
x = torch.randn(1, 128).cuda()
return [x]

def get_init_inputs():
return []



The example new architecture with a custom CUDA kernel looks like this:

python
import torch
from torch.utils.cpp_extension import load_inline

relu_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

global void relu_kernel(const float* x, float* y, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
y[idx] = fmaxf(x[idx], 0.f);
}
}

torch::Tensor relu_cuda(torch::Tensor x) {
auto size = x.numel();
auto y = torch::empty_like(x);
const int block_size = 256;
int num_blocks = (size + block_size - 1) / block_size;
relu_kernel<<<num_blocks, block_size>>>(x.data_ptr<float>(), y.data_ptr<float>(), size);
return y;
}
"""

relu_cpp_source = """
torch::Tensor relu_cuda(torch::Tensor x);
"""

Compile the inline CUDA code
relu = load_inline(
name=“relu”,
cpp_sources=relu_cpp_source,
cuda_sources=relu_source,
functions=[“relu_cuda”],
verbose=True
)

class ModelNew(torch.nn.Module):
def init(self):
super(ModelNew, self).init()
self.relu = relu # The module containing the kernel

def forward(self, x):
    return self.relu.relu_cuda(x)
def get_inputs():
x = torch.randn(1, 128).cuda()
return [x]

def get_init_inputs():
return []



---

Now, you are given the following PyTorch architecture to accelerate. The model computes the Minkowski distance between two tensors and then applies a ReLU activation to the resulting distances. This baseline implementation uses standard element-wise operations to ensure a clear, sample-by-sample calculation.

python
import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
“”"
Minkowski Distance followed by a ReLU activation.
This version uses standard PyTorch operations for a fair baseline.
“”"
def init(self, p=2):
super(Model, self).init()
self.p = p
if p <= 0:
raise ValueError(“p must be positive”)

def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    """
    Compute the ReLU of the Minkowski distance between x and y.
    """
    # Input validation
    if x.shape != y.shape:
        raise ValueError(f"Input tensors must have the same shape, got {x.shape} and {y.shape}")
    if x.dim() != 2:
        raise ValueError(f"Input tensors must be 2D, got {x.dim()}D")
    
    # Step 1: Compute Minkowski distance using element-wise operations
    abs_diff = torch.abs(x - y)

    if self.p == 1:
        d = torch.sum(abs_diff, dim=1)
    elif self.p == 2:
        d = torch.sqrt(torch.sum(abs_diff ** 2, dim=1))
    else:
        d = torch.pow(torch.sum(torch.pow(abs_diff, self.p), dim=1), 1.0/self.p)
    
    # Step 2: Apply ReLU activation
    output = F.relu(d)
    
    return output
batch_size = 256
feature_dim = 512

def get_inputs():
x = torch.randn(batch_size, feature_dim)
y = torch.randn(batch_size, feature_dim)
return [x, y]

def get_init_inputs():
return [2] # p value (default: Euclidean distance)



Your task is to generate the `ModelNew` architecture with a custom CUDA kernel that fuses the Minkowski distance calculation and the ReLU activation into a single kernel launch, thereby eliminating the intermediate distance tensor.

**CRITICAL REQUIREMENTS:**

1.  **Operator Fusion:** The entire logic—computing the Minkowski distance for each sample in the batch and then applying the ReLU activation (`max(distance, 0)`)—must be performed inside a **single CUDA kernel**. No intermediate distance tensors should be written to global memory.
2.  **Algorithmic Specialization:** The implementation must provide specialized, highly optimized kernels for the most common cases, `p=1` (Manhattan) and `p=2` (Euclidean), in addition to a general kernel for any `p`.
3.  **Kernel Logic:**
    *   Each thread block should be responsible for computing the final output for a single sample in the batch.
    *   For `p=1` and `p=2`, the kernel should be a simple loop without shared memory reduction for maximum efficiency.
    *   For the general `p` case, the kernel should use a multi-threaded reduction within a thread block with `extern __shared__`.
4.  **Final Calculation:** Inside the kernel, after the distance is computed by the first thread (`tid == 0`), the ReLU activation must be applied immediately using `fmaxf(distance, 0.0f)` before writing the final result to the output tensor.
5.  **Performance Optimization:** The host-side function should dispatch to the appropriate specialized kernel based on the value of `p` at runtime.
6.  **Code Structure:** Follow the exact structure of the provided example, including `load_inline`, the CUDA source string, the C++ wrapper source string, and the `ModelNew` class. The `get_init_inputs` function must return `[2]` to match the baseline.
7.  **No Fast Math:** Do not use `--use_fast_math` in the compilation flags to ensure numerical accuracy.